<!-- HTML header for doxygen 1.8.9.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<meta name="robots" content="NOINDEX, NOFOLLOW" /> <!-- Prevent indexing by search engines -->
<title>Compute Library: arm_compute/core/CL/ICLKernel.h Source File</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtreedata.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
  $(window).load(resizeHeight);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
  MathJax.Hub.Config({
    extensions: ["tex2jax.js"],
    jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">Compute Library
   &#160;<span id="projectnumber">18.05</span>
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.xhtml"><span>Main&#160;Page</span></a></li>
      <li><a href="pages.xhtml"><span>Related&#160;Pages</span></a></li>
      <li><a href="namespaces.xhtml"><span>Namespaces</span></a></li>
      <li><a href="annotated.xhtml"><span>Data&#160;Structures</span></a></li>
      <li class="current"><a href="files.xhtml"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="files.xhtml"><span>File&#160;List</span></a></li>
      <li><a href="globals.xhtml"><span>Globals</span></a></li>
    </ul>
  </div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
$(document).ready(function(){initNavTree('_i_c_l_kernel_8h_source.xhtml','');});
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">ICLKernel.h</div>  </div>
</div><!--header-->
<div class="contents">
<a href="_i_c_l_kernel_8h.xhtml">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno">    1</span>&#160;<span class="comment">/*</span></div><div class="line"><a name="l00002"></a><span class="lineno">    2</span>&#160;<span class="comment"> * Copyright (c) 2016-2018 ARM Limited.</span></div><div class="line"><a name="l00003"></a><span class="lineno">    3</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno">    4</span>&#160;<span class="comment"> * SPDX-License-Identifier: MIT</span></div><div class="line"><a name="l00005"></a><span class="lineno">    5</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00006"></a><span class="lineno">    6</span>&#160;<span class="comment"> * Permission is hereby granted, free of charge, to any person obtaining a copy</span></div><div class="line"><a name="l00007"></a><span class="lineno">    7</span>&#160;<span class="comment"> * of this software and associated documentation files (the &quot;Software&quot;), to</span></div><div class="line"><a name="l00008"></a><span class="lineno">    8</span>&#160;<span class="comment"> * deal in the Software without restriction, including without limitation the</span></div><div class="line"><a name="l00009"></a><span class="lineno">    9</span>&#160;<span class="comment"> * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or</span></div><div class="line"><a name="l00010"></a><span class="lineno">   10</span>&#160;<span class="comment"> * sell copies of the Software, and to permit persons to whom the Software is</span></div><div class="line"><a name="l00011"></a><span class="lineno">   11</span>&#160;<span class="comment"> * furnished to do so, subject to the following conditions:</span></div><div class="line"><a name="l00012"></a><span class="lineno">   12</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00013"></a><span class="lineno">   13</span>&#160;<span class="comment"> * The above copyright notice and this permission notice shall be included in all</span></div><div class="line"><a name="l00014"></a><span class="lineno">   14</span>&#160;<span class="comment"> * copies or substantial portions of the Software.</span></div><div class="line"><a name="l00015"></a><span class="lineno">   15</span>&#160;<span class="comment"> *</span></div><div class="line"><a name="l00016"></a><span class="lineno">   16</span>&#160;<span class="comment"> * THE SOFTWARE IS PROVIDED &quot;AS IS&quot;, WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</span></div><div class="line"><a name="l00017"></a><span class="lineno">   17</span>&#160;<span class="comment"> * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,</span></div><div class="line"><a name="l00018"></a><span class="lineno">   18</span>&#160;<span class="comment"> * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE</span></div><div class="line"><a name="l00019"></a><span class="lineno">   19</span>&#160;<span class="comment"> * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER</span></div><div class="line"><a name="l00020"></a><span class="lineno">   20</span>&#160;<span class="comment"> * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,</span></div><div class="line"><a name="l00021"></a><span class="lineno">   21</span>&#160;<span class="comment"> * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE</span></div><div class="line"><a name="l00022"></a><span class="lineno">   22</span>&#160;<span class="comment"> * SOFTWARE.</span></div><div class="line"><a name="l00023"></a><span class="lineno">   23</span>&#160;<span class="comment"> */</span></div><div class="line"><a name="l00024"></a><span class="lineno">   24</span>&#160;<span class="preprocessor">#ifndef __ARM_COMPUTE_ICLKERNEL_H__</span></div><div class="line"><a name="l00025"></a><span class="lineno">   25</span>&#160;<span class="preprocessor">#define __ARM_COMPUTE_ICLKERNEL_H__</span></div><div class="line"><a name="l00026"></a><span class="lineno">   26</span>&#160;</div><div class="line"><a name="l00027"></a><span class="lineno">   27</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_c_l_kernel_library_8h.xhtml">arm_compute/core/CL/CLKernelLibrary.h</a>&quot;</span></div><div class="line"><a name="l00028"></a><span class="lineno">   28</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_c_l_types_8h.xhtml">arm_compute/core/CL/CLTypes.h</a>&quot;</span></div><div class="line"><a name="l00029"></a><span class="lineno">   29</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_open_c_l_8h.xhtml">arm_compute/core/CL/OpenCL.h</a>&quot;</span></div><div class="line"><a name="l00030"></a><span class="lineno">   30</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_g_p_u_target_8h.xhtml">arm_compute/core/GPUTarget.h</a>&quot;</span></div><div class="line"><a name="l00031"></a><span class="lineno">   31</span>&#160;<span class="preprocessor">#include &quot;<a class="code" href="_i_kernel_8h.xhtml">arm_compute/core/IKernel.h</a>&quot;</span></div><div class="line"><a name="l00032"></a><span class="lineno">   32</span>&#160;</div><div class="line"><a name="l00033"></a><span class="lineno">   33</span>&#160;<span class="preprocessor">#include &lt;string&gt;</span></div><div class="line"><a name="l00034"></a><span class="lineno">   34</span>&#160;</div><div class="line"><a name="l00035"></a><span class="lineno">   35</span>&#160;<span class="keyword">namespace </span><a class="code" href="namespacearm__compute.xhtml">arm_compute</a></div><div class="line"><a name="l00036"></a><span class="lineno">   36</span>&#160;{</div><div class="line"><a name="l00037"></a><span class="lineno">   37</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div><div class="line"><a name="l00038"></a><span class="lineno">   38</span>&#160;<span class="keyword">class </span>ICLArray;</div><div class="line"><a name="l00039"></a><span class="lineno">   39</span>&#160;<span class="keyword">class </span>ICLTensor;</div><div class="line"><a name="l00040"></a><span class="lineno">   40</span>&#160;<span class="keyword">class </span>Window;</div><div class="line"><a name="l00041"></a><span class="lineno">   41</span>&#160;</div><div class="line"><a name="l00043"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml">   43</a></span>&#160;<span class="keyword">class </span><a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml">ICLKernel</a> : <span class="keyword">public</span> <a class="code" href="classarm__compute_1_1_i_kernel.xhtml">IKernel</a></div><div class="line"><a name="l00044"></a><span class="lineno">   44</span>&#160;{</div><div class="line"><a name="l00045"></a><span class="lineno">   45</span>&#160;<span class="keyword">private</span>:</div><div class="line"><a name="l00050"></a><span class="lineno">   50</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>        dimension_size&gt;</div><div class="line"><a name="l00051"></a><span class="lineno">   51</span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_arguments_per_array()</div><div class="line"><a name="l00052"></a><span class="lineno">   52</span>&#160;    {</div><div class="line"><a name="l00053"></a><span class="lineno">   53</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;dimension_size&gt;();</div><div class="line"><a name="l00054"></a><span class="lineno">   54</span>&#160;    }</div><div class="line"><a name="l00059"></a><span class="lineno">   59</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span>        dimension_size&gt;</div><div class="line"><a name="l00060"></a><span class="lineno">   60</span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_arguments_per_tensor()</div><div class="line"><a name="l00061"></a><span class="lineno">   61</span>&#160;    {</div><div class="line"><a name="l00062"></a><span class="lineno">   62</span>&#160;        <span class="keywordflow">return</span> 2 + 2 * dimension_size;</div><div class="line"><a name="l00063"></a><span class="lineno">   63</span>&#160;    }</div><div class="line"><a name="l00064"></a><span class="lineno">   64</span>&#160;</div><div class="line"><a name="l00065"></a><span class="lineno">   65</span>&#160;<span class="keyword">public</span>:</div><div class="line"><a name="l00067"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">   67</a></span>&#160;    <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">ICLKernel</a>()</div><div class="line"><a name="l00068"></a><span class="lineno">   68</span>&#160;        : _kernel(nullptr), _lws_hint(<a class="code" href="classarm__compute_1_1_c_l_kernel_library.xhtml">CLKernelLibrary</a>::get().default_ndrange()), _target(<a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a>::<a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3afa362f4daec88442a387ff7cda411a38">MIDGARD</a>), _config_id(<a class="code" href="namespacearm__compute.xhtml">arm_compute</a>::default_config_id), _max_workgroup_size(0)</div><div class="line"><a name="l00069"></a><span class="lineno">   69</span>&#160;    {</div><div class="line"><a name="l00070"></a><span class="lineno">   70</span>&#160;    }</div><div class="line"><a name="l00075"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">   75</a></span>&#160;    cl::Kernel &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">kernel</a>()</div><div class="line"><a name="l00076"></a><span class="lineno">   76</span>&#160;    {</div><div class="line"><a name="l00077"></a><span class="lineno">   77</span>&#160;        <span class="keywordflow">return</span> _kernel;</div><div class="line"><a name="l00078"></a><span class="lineno">   78</span>&#160;    }</div><div class="line"><a name="l00087"></a><span class="lineno">   87</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div><div class="line"><a name="l00088"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">   88</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">add_1D_array_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00089"></a><span class="lineno">   89</span>&#160;    {</div><div class="line"><a name="l00090"></a><span class="lineno">   90</span>&#160;        add_array_argument&lt;T, 1&gt;(idx, array, strides, num_dimensions, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00091"></a><span class="lineno">   91</span>&#160;    }</div><div class="line"><a name="l00098"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">   98</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">add_1D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00099"></a><span class="lineno">   99</span>&#160;    {</div><div class="line"><a name="l00100"></a><span class="lineno">  100</span>&#160;        add_tensor_argument&lt;1&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00101"></a><span class="lineno">  101</span>&#160;    }</div><div class="line"><a name="l00108"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">  108</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">add_2D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00109"></a><span class="lineno">  109</span>&#160;    {</div><div class="line"><a name="l00110"></a><span class="lineno">  110</span>&#160;        add_tensor_argument&lt;2&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00111"></a><span class="lineno">  111</span>&#160;    }</div><div class="line"><a name="l00118"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">  118</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">add_3D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00119"></a><span class="lineno">  119</span>&#160;    {</div><div class="line"><a name="l00120"></a><span class="lineno">  120</span>&#160;        add_tensor_argument&lt;3&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00121"></a><span class="lineno">  121</span>&#160;    }</div><div class="line"><a name="l00128"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">  128</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">add_4D_tensor_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>)</div><div class="line"><a name="l00129"></a><span class="lineno">  129</span>&#160;    {</div><div class="line"><a name="l00130"></a><span class="lineno">  130</span>&#160;        add_tensor_argument&lt;4&gt;(idx, tensor, <a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>);</div><div class="line"><a name="l00131"></a><span class="lineno">  131</span>&#160;    }</div><div class="line"><a name="l00136"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a278f0e6c68ca17e71f4c4ff82f360aa0">  136</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a278f0e6c68ca17e71f4c4ff82f360aa0">num_arguments_per_1D_array</a>()</div><div class="line"><a name="l00137"></a><span class="lineno">  137</span>&#160;    {</div><div class="line"><a name="l00138"></a><span class="lineno">  138</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_array&lt;1&gt;();</div><div class="line"><a name="l00139"></a><span class="lineno">  139</span>&#160;    }</div><div class="line"><a name="l00144"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a43b6c5e4b57069c5f61e96dff24c212d">  144</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a43b6c5e4b57069c5f61e96dff24c212d">num_arguments_per_1D_tensor</a>()</div><div class="line"><a name="l00145"></a><span class="lineno">  145</span>&#160;    {</div><div class="line"><a name="l00146"></a><span class="lineno">  146</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;1&gt;();</div><div class="line"><a name="l00147"></a><span class="lineno">  147</span>&#160;    }</div><div class="line"><a name="l00152"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a45601e0d46621a5b6f2e417d60e5c800">  152</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a45601e0d46621a5b6f2e417d60e5c800">num_arguments_per_2D_tensor</a>()</div><div class="line"><a name="l00153"></a><span class="lineno">  153</span>&#160;    {</div><div class="line"><a name="l00154"></a><span class="lineno">  154</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;2&gt;();</div><div class="line"><a name="l00155"></a><span class="lineno">  155</span>&#160;    }</div><div class="line"><a name="l00160"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6c9c1e7a7d96743375ca40847f0f12e2">  160</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6c9c1e7a7d96743375ca40847f0f12e2">num_arguments_per_3D_tensor</a>()</div><div class="line"><a name="l00161"></a><span class="lineno">  161</span>&#160;    {</div><div class="line"><a name="l00162"></a><span class="lineno">  162</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;3&gt;();</div><div class="line"><a name="l00163"></a><span class="lineno">  163</span>&#160;    }</div><div class="line"><a name="l00168"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a184fdf37587a9314cf12623accea6c73">  168</a></span>&#160;    constexpr <span class="keyword">static</span> <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a184fdf37587a9314cf12623accea6c73">num_arguments_per_4D_tensor</a>()</div><div class="line"><a name="l00169"></a><span class="lineno">  169</span>&#160;    {</div><div class="line"><a name="l00170"></a><span class="lineno">  170</span>&#160;        <span class="keywordflow">return</span> num_arguments_per_tensor&lt;4&gt;();</div><div class="line"><a name="l00171"></a><span class="lineno">  171</span>&#160;    }</div><div class="line"><a name="l00179"></a><span class="lineno">  179</span>&#160;    <span class="keyword">virtual</span> <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#af6a174d47571f51f199ffc27ecc10f51">run</a>(<span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>, cl::CommandQueue &amp;queue) = 0;</div><div class="line"><a name="l00185"></a><span class="lineno">  185</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div><div class="line"><a name="l00186"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">  186</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">add_argument</a>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, T value)</div><div class="line"><a name="l00187"></a><span class="lineno">  187</span>&#160;    {</div><div class="line"><a name="l00188"></a><span class="lineno">  188</span>&#160;        _kernel.setArg(idx++, value);</div><div class="line"><a name="l00189"></a><span class="lineno">  189</span>&#160;    }</div><div class="line"><a name="l00190"></a><span class="lineno">  190</span>&#160;</div><div class="line"><a name="l00197"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad356b88c8f61267d593d9ed99835bde9">  197</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad356b88c8f61267d593d9ed99835bde9">set_lws_hint</a>(<span class="keyword">const</span> cl::NDRange &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a>)</div><div class="line"><a name="l00198"></a><span class="lineno">  198</span>&#160;    {</div><div class="line"><a name="l00199"></a><span class="lineno">  199</span>&#160;        _lws_hint = <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a>;</div><div class="line"><a name="l00200"></a><span class="lineno">  200</span>&#160;    }</div><div class="line"><a name="l00201"></a><span class="lineno">  201</span>&#160;</div><div class="line"><a name="l00206"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">  206</a></span>&#160;    cl::NDRange <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a>()<span class="keyword"> const</span></div><div class="line"><a name="l00207"></a><span class="lineno">  207</span>&#160;<span class="keyword">    </span>{</div><div class="line"><a name="l00208"></a><span class="lineno">  208</span>&#160;        <span class="keywordflow">return</span> _lws_hint;</div><div class="line"><a name="l00209"></a><span class="lineno">  209</span>&#160;    }</div><div class="line"><a name="l00210"></a><span class="lineno">  210</span>&#160;</div><div class="line"><a name="l00221"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">  221</a></span>&#160;    <span class="keyword">const</span> std::string &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">config_id</a>()<span class="keyword"> const</span></div><div class="line"><a name="l00222"></a><span class="lineno">  222</span>&#160;<span class="keyword">    </span>{</div><div class="line"><a name="l00223"></a><span class="lineno">  223</span>&#160;        <span class="keywordflow">return</span> _config_id;</div><div class="line"><a name="l00224"></a><span class="lineno">  224</span>&#160;    }</div><div class="line"><a name="l00225"></a><span class="lineno">  225</span>&#160;</div><div class="line"><a name="l00230"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">  230</a></span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">set_target</a>(<a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a> target)</div><div class="line"><a name="l00231"></a><span class="lineno">  231</span>&#160;    {</div><div class="line"><a name="l00232"></a><span class="lineno">  232</span>&#160;        _target = target;</div><div class="line"><a name="l00233"></a><span class="lineno">  233</span>&#160;    }</div><div class="line"><a name="l00234"></a><span class="lineno">  234</span>&#160;</div><div class="line"><a name="l00239"></a><span class="lineno">  239</span>&#160;    <span class="keywordtype">void</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">set_target</a>(cl::Device &amp;device);</div><div class="line"><a name="l00240"></a><span class="lineno">  240</span>&#160;</div><div class="line"><a name="l00245"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">  245</a></span>&#160;    <a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">get_target</a>()<span class="keyword"> const</span></div><div class="line"><a name="l00246"></a><span class="lineno">  246</span>&#160;<span class="keyword">    </span>{</div><div class="line"><a name="l00247"></a><span class="lineno">  247</span>&#160;        <span class="keywordflow">return</span> _target;</div><div class="line"><a name="l00248"></a><span class="lineno">  248</span>&#160;    }</div><div class="line"><a name="l00249"></a><span class="lineno">  249</span>&#160;</div><div class="line"><a name="l00254"></a><span class="lineno">  254</span>&#160;    <span class="keywordtype">size_t</span> <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#abca336f832d730e8494049bd714df60a">get_max_workgroup_size</a>();</div><div class="line"><a name="l00261"></a><span class="lineno">  261</span>&#160;    <span class="keyword">static</span> cl::NDRange <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a4b4a20e518c04e657839280cd486a85d">gws_from_window</a>(<span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div><div class="line"><a name="l00262"></a><span class="lineno">  262</span>&#160;</div><div class="line"><a name="l00263"></a><span class="lineno">  263</span>&#160;<span class="keyword">private</span>:</div><div class="line"><a name="l00272"></a><span class="lineno">  272</span>&#160;    <span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div><div class="line"><a name="l00273"></a><span class="lineno">  273</span>&#160;    <span class="keywordtype">void</span> add_array_argument(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div><div class="line"><a name="l00280"></a><span class="lineno">  280</span>&#160;    <span class="keyword">template</span> &lt;<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div><div class="line"><a name="l00281"></a><span class="lineno">  281</span>&#160;    <span class="keywordtype">void</span> add_tensor_argument(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_tensor.xhtml">ICLTensor</a> *tensor, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window);</div><div class="line"><a name="l00282"></a><span class="lineno">  282</span>&#160;</div><div class="line"><a name="l00283"></a><span class="lineno">  283</span>&#160;<span class="keyword">protected</span>:</div><div class="line"><a name="l00284"></a><span class="lineno">  284</span>&#160;    cl::Kernel  _kernel;             </div><div class="line"><a name="l00285"></a><span class="lineno">  285</span>&#160;    cl::NDRange _lws_hint;           </div><div class="line"><a name="l00286"></a><span class="lineno">  286</span>&#160;    <a class="code" href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">GPUTarget</a>   _target;             </div><div class="line"><a name="l00287"></a><span class="lineno">  287</span>&#160;    std::string _config_id;          </div><div class="line"><a name="l00288"></a><span class="lineno">  288</span>&#160;    <span class="keywordtype">size_t</span>      _max_workgroup_size; </div><div class="line"><a name="l00289"></a><span class="lineno">  289</span>&#160;};</div><div class="line"><a name="l00290"></a><span class="lineno">  290</span>&#160;</div><div class="line"><a name="l00304"></a><span class="lineno">  304</span>&#160;<span class="keywordtype">void</span> <a class="code" href="namespacearm__compute.xhtml#a0f68015f649819f1102aa73d7b88b29e">enqueue</a>(cl::CommandQueue &amp;queue, <a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml">ICLKernel</a> &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">kernel</a>, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;<a class="code" href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">window</a>, <span class="keyword">const</span> cl::NDRange &amp;<a class="code" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">lws_hint</a> = <a class="code" href="classarm__compute_1_1_c_l_kernel_library.xhtml#aebe2d25607545d31f2b1ace6d60a81da">CLKernelLibrary::get</a>().default_ndrange());</div><div class="line"><a name="l00305"></a><span class="lineno">  305</span>&#160;</div><div class="line"><a name="l00314"></a><span class="lineno">  314</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension_size&gt;</div><div class="line"><a name="l00315"></a><span class="lineno"><a class="line" href="classarm__compute_1_1_i_c_l_kernel.xhtml#a2d7c6b5f3332604ad6a637457f65c342">  315</a></span>&#160;<span class="keywordtype">void</span> ICLKernel::add_array_argument(<span class="keywordtype">unsigned</span> &amp;idx, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml">ICLArray&lt;T&gt;</a> *array, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_strides.xhtml">Strides</a> &amp;strides, <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> num_dimensions, <span class="keyword">const</span> <a class="code" href="classarm__compute_1_1_window.xhtml">Window</a> &amp;window)</div><div class="line"><a name="l00316"></a><span class="lineno">  316</span>&#160;{</div><div class="line"><a name="l00317"></a><span class="lineno">  317</span>&#160;    <a class="code" href="_error_8h.xhtml#a54a6080c9f4df1f908e57a9bbb46f5da">ARM_COMPUTE_ERROR_ON</a>(array == <span class="keyword">nullptr</span>);</div><div class="line"><a name="l00318"></a><span class="lineno">  318</span>&#160;</div><div class="line"><a name="l00319"></a><span class="lineno">  319</span>&#160;    <span class="comment">// Calculate offset to the start of the window</span></div><div class="line"><a name="l00320"></a><span class="lineno">  320</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> offset_first_element = 0;</div><div class="line"><a name="l00321"></a><span class="lineno">  321</span>&#160;</div><div class="line"><a name="l00322"></a><span class="lineno">  322</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> n = 0; n &lt; num_dimensions; ++n)</div><div class="line"><a name="l00323"></a><span class="lineno">  323</span>&#160;    {</div><div class="line"><a name="l00324"></a><span class="lineno">  324</span>&#160;        offset_first_element += window[n].start() * strides[n];</div><div class="line"><a name="l00325"></a><span class="lineno">  325</span>&#160;    }</div><div class="line"><a name="l00326"></a><span class="lineno">  326</span>&#160;</div><div class="line"><a name="l00327"></a><span class="lineno">  327</span>&#160;    <span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> idx_start = idx;</div><div class="line"><a name="l00328"></a><span class="lineno">  328</span>&#160;    _kernel.setArg(idx++, array-&gt;<a class="code" href="classarm__compute_1_1_i_c_l_array.xhtml#a1fb4c50755a0ef424652246838ed91a6">cl_buffer</a>());</div><div class="line"><a name="l00329"></a><span class="lineno">  329</span>&#160;</div><div class="line"><a name="l00330"></a><span class="lineno">  330</span>&#160;    <span class="keywordflow">for</span>(<span class="keywordtype">unsigned</span> <span class="keywordtype">int</span> dimension = 0; dimension &lt; dimension_size; dimension++)</div><div class="line"><a name="l00331"></a><span class="lineno">  331</span>&#160;    {</div><div class="line"><a name="l00332"></a><span class="lineno">  332</span>&#160;        _kernel.setArg&lt;cl_uint&gt;(idx++, strides[dimension]);</div><div class="line"><a name="l00333"></a><span class="lineno">  333</span>&#160;        _kernel.setArg&lt;cl_uint&gt;(idx++, strides[dimension] * window[dimension].step());</div><div class="line"><a name="l00334"></a><span class="lineno">  334</span>&#160;    }</div><div class="line"><a name="l00335"></a><span class="lineno">  335</span>&#160;</div><div class="line"><a name="l00336"></a><span class="lineno">  336</span>&#160;    _kernel.setArg&lt;cl_uint&gt;(idx++, offset_first_element);</div><div class="line"><a name="l00337"></a><span class="lineno">  337</span>&#160;</div><div class="line"><a name="l00338"></a><span class="lineno">  338</span>&#160;    <a class="code" href="_error_8h.xhtml#a5bbdcf574d3f5e412fa6a1117911e67b">ARM_COMPUTE_ERROR_ON_MSG</a>(idx_start + num_arguments_per_array&lt;dimension_size&gt;() != idx,</div><div class="line"><a name="l00339"></a><span class="lineno">  339</span>&#160;                             <span class="stringliteral">&quot;add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel&quot;</span>, dimension_size, num_arguments_per_array&lt;dimension_size&gt;());</div><div class="line"><a name="l00340"></a><span class="lineno">  340</span>&#160;    <a class="code" href="_error_8h.xhtml#a6dc630a6ae9cc063b3924bcea8dee9d6">ARM_COMPUTE_UNUSED</a>(idx_start);</div><div class="line"><a name="l00341"></a><span class="lineno">  341</span>&#160;}</div><div class="line"><a name="l00342"></a><span class="lineno">  342</span>&#160;}</div><div class="line"><a name="l00343"></a><span class="lineno">  343</span>&#160;<span class="preprocessor">#endif </span><span class="comment">/*__ARM_COMPUTE_ICLKERNEL_H__ */</span><span class="preprocessor"></span></div><div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a43b6c5e4b57069c5f61e96dff24c212d"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a43b6c5e4b57069c5f61e96dff24c212d">arm_compute::ICLKernel::num_arguments_per_1D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_1D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 1D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00144">ICLKernel.h:144</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_kernel_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_kernel.xhtml">arm_compute::IKernel</a></div><div class="ttdoc">Common information for all the kernels. </div><div class="ttdef"><b>Definition:</b> <a href="_i_kernel_8h_source.xhtml#l00033">IKernel.h:33</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a278f0e6c68ca17e71f4c4ff82f360aa0"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a278f0e6c68ca17e71f4c4ff82f360aa0">arm_compute::ICLKernel::num_arguments_per_1D_array</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_1D_array()</div><div class="ttdoc">Returns the number of arguments enqueued per 1D array object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00136">ICLKernel.h:136</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a0737b867334392d5dc5c2e3291d67dbf"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a0737b867334392d5dc5c2e3291d67dbf">arm_compute::ICLKernel::lws_hint</a></div><div class="ttdeci">cl::NDRange lws_hint() const </div><div class="ttdoc">Return the Local-Workgroup-Size hint. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00206">ICLKernel.h:206</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ae5121015ab09ece4d470f50c7ffe198e"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ae5121015ab09ece4d470f50c7ffe198e">arm_compute::ICLKernel::kernel</a></div><div class="ttdeci">cl::Kernel &amp; kernel()</div><div class="ttdoc">Returns a reference to the OpenCL kernel of this object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00075">ICLKernel.h:75</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ad356b88c8f61267d593d9ed99835bde9"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad356b88c8f61267d593d9ed99835bde9">arm_compute::ICLKernel::set_lws_hint</a></div><div class="ttdeci">void set_lws_hint(const cl::NDRange &amp;lws_hint)</div><div class="ttdoc">Set the Local-Workgroup-Size hint. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00197">ICLKernel.h:197</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a50f427a1d9419800972b9e03c4034311"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a50f427a1d9419800972b9e03c4034311">arm_compute::ICLKernel::add_argument</a></div><div class="ttdeci">void add_argument(unsigned int &amp;idx, T value)</div><div class="ttdoc">Add the passed parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00186">ICLKernel.h:186</a></div></div>
<div class="ttc" id="_c_l_types_8h_xhtml"><div class="ttname"><a href="_c_l_types_8h.xhtml">CLTypes.h</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a9331d385192a50adf74d3af40ce0fa20"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a9331d385192a50adf74d3af40ce0fa20">arm_compute::ICLKernel::add_1D_array_argument</a></div><div class="ttdeci">void add_1D_array_argument(unsigned int &amp;idx, const ICLArray&lt; T &gt; *array, const Strides &amp;strides, unsigned int num_dimensions, const Window &amp;window)</div><div class="ttdoc">Add the passed 1D array&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00088">ICLKernel.h:88</a></div></div>
<div class="ttc" id="_error_8h_xhtml_a54a6080c9f4df1f908e57a9bbb46f5da"><div class="ttname"><a href="_error_8h.xhtml#a54a6080c9f4df1f908e57a9bbb46f5da">ARM_COMPUTE_ERROR_ON</a></div><div class="ttdeci">#define ARM_COMPUTE_ERROR_ON(cond)</div><div class="ttdoc">If the condition is true then an error message is printed and an exception thrown. </div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00328">Error.h:328</a></div></div>
<div class="ttc" id="_i_kernel_8h_xhtml"><div class="ttname"><a href="_i_kernel_8h.xhtml">IKernel.h</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml">arm_compute::ICLKernel</a></div><div class="ttdoc">Common interface for all the OpenCL kernels. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00043">ICLKernel.h:43</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a28f5847162f352444c6ac1825d0e99c7"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a28f5847162f352444c6ac1825d0e99c7">arm_compute::ICLKernel::add_3D_tensor_argument</a></div><div class="ttdeci">void add_3D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 3D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00118">ICLKernel.h:118</a></div></div>
<div class="ttc" id="namespacearm__compute_xhtml"><div class="ttname"><a href="namespacearm__compute.xhtml">arm_compute</a></div><div class="ttdoc">This file contains all available output stages for GEMMLowp on OpenCL. </div><div class="ttdef"><b>Definition:</b> <a href="00__introduction_8dox_source.xhtml#l00001">00_introduction.dox:1</a></div></div>
<div class="ttc" id="classarm__compute_1_1_c_l_kernel_library_xhtml_aebe2d25607545d31f2b1ace6d60a81da"><div class="ttname"><a href="classarm__compute_1_1_c_l_kernel_library.xhtml#aebe2d25607545d31f2b1ace6d60a81da">arm_compute::CLKernelLibrary::get</a></div><div class="ttdeci">static CLKernelLibrary &amp; get()</div><div class="ttdoc">Access the KernelLibrary singleton. </div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_array_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_array.xhtml">arm_compute::ICLArray</a></div><div class="ttdoc">Interface for OpenCL Array. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_array_8h_source.xhtml#l00035">ICLArray.h:35</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a6c9c1e7a7d96743375ca40847f0f12e2"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6c9c1e7a7d96743375ca40847f0f12e2">arm_compute::ICLKernel::num_arguments_per_3D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_3D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 3D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00160">ICLKernel.h:160</a></div></div>
<div class="ttc" id="_error_8h_xhtml_a6dc630a6ae9cc063b3924bcea8dee9d6"><div class="ttname"><a href="_error_8h.xhtml#a6dc630a6ae9cc063b3924bcea8dee9d6">ARM_COMPUTE_UNUSED</a></div><div class="ttdeci">#define ARM_COMPUTE_UNUSED(...)</div><div class="ttdoc">To avoid unused variables warnings. </div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00159">Error.h:159</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a646cd535a16835b246c3367a63d96250"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a646cd535a16835b246c3367a63d96250">arm_compute::ICLKernel::get_target</a></div><div class="ttdeci">GPUTarget get_target() const </div><div class="ttdoc">Get the targeted GPU architecture. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00245">ICLKernel.h:245</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a45601e0d46621a5b6f2e417d60e5c800"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a45601e0d46621a5b6f2e417d60e5c800">arm_compute::ICLKernel::num_arguments_per_2D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_2D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 2D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00152">ICLKernel.h:152</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a184fdf37587a9314cf12623accea6c73"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a184fdf37587a9314cf12623accea6c73">arm_compute::ICLKernel::num_arguments_per_4D_tensor</a></div><div class="ttdeci">static constexpr unsigned int num_arguments_per_4D_tensor()</div><div class="ttdoc">Returns the number of arguments enqueued per 4D tensor object. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00168">ICLKernel.h:168</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a226d2e9e6d3c42d681666566fe950b2d"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a226d2e9e6d3c42d681666566fe950b2d">arm_compute::ICLKernel::config_id</a></div><div class="ttdeci">const std::string &amp; config_id() const </div><div class="ttdoc">Get the configuration ID. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00221">ICLKernel.h:221</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a4b4a20e518c04e657839280cd486a85d"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a4b4a20e518c04e657839280cd486a85d">arm_compute::ICLKernel::gws_from_window</a></div><div class="ttdeci">static cl::NDRange gws_from_window(const Window &amp;window)</div><div class="ttdoc">Get the global work size given an execution window. </div></div>
<div class="ttc" id="classarm__compute_1_1_strides_xhtml"><div class="ttname"><a href="classarm__compute_1_1_strides.xhtml">arm_compute::Strides</a></div><div class="ttdoc">Strides of an item in bytes. </div><div class="ttdef"><b>Definition:</b> <a href="_strides_8h_source.xhtml#l00037">Strides.h:37</a></div></div>
<div class="ttc" id="namespacearm__compute_xhtml_a0f68015f649819f1102aa73d7b88b29e"><div class="ttname"><a href="namespacearm__compute.xhtml#a0f68015f649819f1102aa73d7b88b29e">arm_compute::enqueue</a></div><div class="ttdeci">void enqueue(cl::CommandQueue &amp;queue, ICLKernel &amp;kernel, const Window &amp;window, const cl::NDRange &amp;lws_hint=CLKernelLibrary::get().default_ndrange())</div><div class="ttdoc">Add the kernel to the command queue with the given window. </div></div>
<div class="ttc" id="_g_p_u_target_8h_xhtml"><div class="ttname"><a href="_g_p_u_target_8h.xhtml">GPUTarget.h</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_af6a174d47571f51f199ffc27ecc10f51"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#af6a174d47571f51f199ffc27ecc10f51">arm_compute::ICLKernel::run</a></div><div class="ttdeci">virtual void run(const Window &amp;window, cl::CommandQueue &amp;queue)=0</div><div class="ttdoc">Enqueue the OpenCL kernel to process the given window on the passed OpenCL command queue...</div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ac74dad3e61f79334f5e73f3c3ac603cb"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ac74dad3e61f79334f5e73f3c3ac603cb">arm_compute::ICLKernel::add_2D_tensor_argument</a></div><div class="ttdeci">void add_2D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 2D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00108">ICLKernel.h:108</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_tensor_xhtml"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_tensor.xhtml">arm_compute::ICLTensor</a></div><div class="ttdoc">Interface for OpenCL tensor. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_tensor_8h_source.xhtml#l00040">ICLTensor.h:40</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a6b10e96ce90bf901d17def86b874b019"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a6b10e96ce90bf901d17def86b874b019">arm_compute::ICLKernel::ICLKernel</a></div><div class="ttdeci">ICLKernel()</div><div class="ttdoc">Constructor. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00067">ICLKernel.h:67</a></div></div>
<div class="ttc" id="namespacearm__compute_xhtml_a735ac6c2a02e320969625308810444f3"><div class="ttname"><a href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3">arm_compute::GPUTarget</a></div><div class="ttdeci">GPUTarget</div><div class="ttdoc">Available GPU Targets. </div><div class="ttdef"><b>Definition:</b> <a href="_g_p_u_target_8h_source.xhtml#l00034">GPUTarget.h:34</a></div></div>
<div class="ttc" id="_c_l_kernel_library_8h_xhtml"><div class="ttname"><a href="_c_l_kernel_library_8h.xhtml">CLKernelLibrary.h</a></div></div>
<div class="ttc" id="namespacearm__compute_xhtml_a735ac6c2a02e320969625308810444f3afa362f4daec88442a387ff7cda411a38"><div class="ttname"><a href="namespacearm__compute.xhtml#a735ac6c2a02e320969625308810444f3afa362f4daec88442a387ff7cda411a38">arm_compute::GPUTarget::MIDGARD</a></div></div>
<div class="ttc" id="_open_c_l_8h_xhtml"><div class="ttname"><a href="_open_c_l_8h.xhtml">OpenCL.h</a></div></div>
<div class="ttc" id="classarm__compute_1_1_c_l_kernel_library_xhtml"><div class="ttname"><a href="classarm__compute_1_1_c_l_kernel_library.xhtml">arm_compute::CLKernelLibrary</a></div><div class="ttdoc">CLKernelLibrary class. </div><div class="ttdef"><b>Definition:</b> <a href="_c_l_kernel_library_8h_source.xhtml#l00192">CLKernelLibrary.h:192</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_abca336f832d730e8494049bd714df60a"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#abca336f832d730e8494049bd714df60a">arm_compute::ICLKernel::get_max_workgroup_size</a></div><div class="ttdeci">size_t get_max_workgroup_size()</div><div class="ttdoc">Get the maximum workgroup size for the device the CLKernelLibrary uses. </div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_ad5ba9d34a3a855bf1dd2e36316ff550a"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#ad5ba9d34a3a855bf1dd2e36316ff550a">arm_compute::ICLKernel::set_target</a></div><div class="ttdeci">void set_target(GPUTarget target)</div><div class="ttdoc">Set the targeted GPU architecture. </div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00230">ICLKernel.h:230</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a479e7043e65dc87de35d374e108510f7"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a479e7043e65dc87de35d374e108510f7">arm_compute::ICLKernel::add_1D_tensor_argument</a></div><div class="ttdeci">void add_1D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 1D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00098">ICLKernel.h:98</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_kernel_xhtml_a33e09c946b338fbfc780a9d1c66e68ad"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_kernel.xhtml#a33e09c946b338fbfc780a9d1c66e68ad">arm_compute::ICLKernel::add_4D_tensor_argument</a></div><div class="ttdeci">void add_4D_tensor_argument(unsigned int &amp;idx, const ICLTensor *tensor, const Window &amp;window)</div><div class="ttdoc">Add the passed 4D tensor&amp;#39;s parameters to the object&amp;#39;s kernel&amp;#39;s arguments starting from the index idx...</div><div class="ttdef"><b>Definition:</b> <a href="_i_c_l_kernel_8h_source.xhtml#l00128">ICLKernel.h:128</a></div></div>
<div class="ttc" id="classarm__compute_1_1_i_c_l_array_xhtml_a1fb4c50755a0ef424652246838ed91a6"><div class="ttname"><a href="classarm__compute_1_1_i_c_l_array.xhtml#a1fb4c50755a0ef424652246838ed91a6">arm_compute::ICLArray::cl_buffer</a></div><div class="ttdeci">virtual const cl::Buffer &amp; cl_buffer() const =0</div><div class="ttdoc">Interface to be implemented by the child class to return a reference to the OpenCL buffer containing ...</div></div>
<div class="ttc" id="classarm__compute_1_1_i_kernel_xhtml_a3f5646133956f06348b310ccc3d36353"><div class="ttname"><a href="classarm__compute_1_1_i_kernel.xhtml#a3f5646133956f06348b310ccc3d36353">arm_compute::IKernel::window</a></div><div class="ttdeci">const Window &amp; window() const </div><div class="ttdoc">The maximum window the kernel can be executed on. </div></div>
<div class="ttc" id="classarm__compute_1_1_window_xhtml"><div class="ttname"><a href="classarm__compute_1_1_window.xhtml">arm_compute::Window</a></div><div class="ttdoc">Describe a multidimensional execution window. </div><div class="ttdef"><b>Definition:</b> <a href="_window_8h_source.xhtml#l00039">Window.h:39</a></div></div>
<div class="ttc" id="_error_8h_xhtml_a5bbdcf574d3f5e412fa6a1117911e67b"><div class="ttname"><a href="_error_8h.xhtml#a5bbdcf574d3f5e412fa6a1117911e67b">ARM_COMPUTE_ERROR_ON_MSG</a></div><div class="ttdeci">#define ARM_COMPUTE_ERROR_ON_MSG(cond,...)</div><div class="ttdef"><b>Definition:</b> <a href="_error_8h_source.xhtml#l00319">Error.h:319</a></div></div>
</div><!-- fragment --></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="dir_214608ef36d61442cb2b0c1c4e9a7def.xhtml">arm_compute</a></li><li class="navelem"><a class="el" href="dir_1fb090f0c6070330bfaccc4236d3ca0d.xhtml">core</a></li><li class="navelem"><a class="el" href="dir_1633709b037eda5b8477677909f78129.xhtml">CL</a></li><li class="navelem"><a class="el" href="_i_c_l_kernel_8h.xhtml">ICLKernel.h</a></li>
    <li class="footer">Generated on Wed May 23 2018 11:36:34 for Compute Library by
    <a href="http://www.doxygen.org/index.html">
    <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.8.11 </li>
  </ul>
</div>
</body>
</html>
